#include <stdio.h>
#include <stdlib.h>
#include <CL/cl.h>
#include <assert.h>

int main() {
  
#define TAILLE_MAT 4
#define TAILLE_BLOC 2
  
  int taille_mat = TAILLE_MAT;
  int taille_bloc = TAILLE_BLOC;

  size_t size = TAILLE_MAT * TAILLE_MAT * sizeof(float);
  
  cl_event ev_writeA, ev_ker1, ev_ker2, ev_ker3;
  cl_event ev_readA;
  int i;
  
  //float * matA = malloc(N * N * sizeof(float));
  //float * matB = malloc(N * N * sizeof(float));
  //float * matC = malloc(N * N * sizeof(float));
  
  //for (i=0; i<N*N; i++) {
  
  float matA[] = {1,1,1,1,
		  1,5,5,5,
		  1,5,14,14,
		  1,5,14,15};
  //matB[i] = (float)i;
  //}
  //for (i=0; i<N*N; i++) {
  //  matB[i] = 0.0;
  // }

   
  cl_uint nb_platf;
  clGetPlatformIDs(0, NULL, &nb_platf);

  printf("Nombre de plateformes: %d\n", nb_platf);

  cl_platform_id platfs[nb_platf];
  clGetPlatformIDs(nb_platf, platfs, NULL);

  size_t plat_name_size;
  clGetPlatformInfo(platfs[0], CL_PLATFORM_NAME, 0, NULL, &plat_name_size);
  char plat_name[plat_name_size];
  clGetPlatformInfo(platfs[0], CL_PLATFORM_NAME, plat_name_size, &plat_name, NULL);
  printf("Nom de la plateforme: %s\n", plat_name);

  size_t plat_vendor_size;
  clGetPlatformInfo(platfs[0], CL_PLATFORM_VENDOR, 0, NULL, &plat_vendor_size);
  char plat_vendor[plat_vendor_size];
  clGetPlatformInfo(platfs[0], CL_PLATFORM_VENDOR, plat_vendor_size, &plat_vendor, NULL);
  printf("Nom de la plateforme: %s\n", plat_vendor);

  cl_uint nb_devs;
  clGetDeviceIDs(platfs[0], CL_DEVICE_TYPE_ALL, 0, NULL, &nb_devs);

  cl_device_id devs[nb_devs];
  clGetDeviceIDs(platfs[0], CL_DEVICE_TYPE_ALL, nb_devs, devs, NULL);

  cl_context ctx = clCreateContext(NULL, nb_devs, devs, NULL, NULL, NULL);
  cl_command_queue cq = clCreateCommandQueue(ctx, devs[0], CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE | CL_QUEUE_PROFILING_ENABLE, NULL);

  FILE * f = fopen("./matrix_cho3k.cl", "r");
  fseek(f, 0, SEEK_END);
  size_t source_size = ftell(f);
  fseek(f, 0, SEEK_SET);

  char * source = malloc(source_size+1);
  fread(source, 1, source_size, f);
  source[source_size] = '\0';
  fclose(f);

  //   printf("Code kernel: %s\n", source);


  //Récuperation du code des kernels.
  cl_program prg1 = clCreateProgramWithSource(ctx, 1, (const char**)&source, NULL, NULL);
  clBuildProgram(prg1, nb_devs, devs, NULL, NULL, NULL);

  size_t log_size;
  clGetProgramBuildInfo(prg1, devs[0], CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size);
  char log[log_size];
  clGetProgramBuildInfo(prg1, devs[0], CL_PROGRAM_BUILD_LOG, log_size, log, NULL);

  printf("Build log: %s\n", log);

  //Boucle sur les blocs
  int indice;
  for(indice=0 ; indice<n/taille_bloc ; indice++)
    {
      //Kernel : Cholesky sur un bloc.
      //1111111111111111111111111111111111111111111111111111111111111111111111111111111111
      cl_kernel ker1 = clCreateKernel(prg1, "matrix_cho1", NULL);

      cl_mem bufA = clCreateBuffer(ctx, CL_MEM_READ_WRITE, size, NULL, NULL);
      //cl_mem bufB = clCreateBuffer(ctx, CL_MEM_READ_ONLY, size, NULL, NULL);
      //cl_mem bufC = clCreateBuffer(ctx, CL_MEM_WRITE_ONLY, size, NULL, NULL);

      clEnqueueWriteBuffer(cq, bufA, 0, 0, size, matA, 0, NULL, &ev_writeA);
      //clEnqueueWriteBuffer(cq, bufB, 0, 0, size, matB, 0, NULL, &ev_writeB);

      clSetKernelArg(ker1, 0, sizeof(bufA), &bufA);
      clSetKernelArg(ker1, 0, sizeof(int), &indice);
      clSetKernelArg(ker2, 0, sizeof(int),&taille_bloc);
      clSetKernelArg(ker1, 0, sizeof(int), &taille_mat);

      //nbre de threads 
      size_t globalDim1[] = {TAILLE_BLOC,TAILLE_BLOC};
      size_t localDim1[] = {TAILLE_BLOC,TAILLE_BLOC};
      cl_event deps1[] = {ev_writeA};
      cl_uint ret1 = clEnqueueNDRangeKernel(cq, ker1, 2, NULL, globalDim1, localDim1, 1, deps1, &ev_ker1);
      printf("erreur : %d\n",ret1);
      //clEnqueueReadBuffer(cq, bufA, 0, 0, size, matA, 1, &ev_ker1, &ev_readA);//Rapatriement de la matrice???? Du coup faut pas!!!!
      
      
      
      
      //Kernel : Mise à jour des éléments sous diagonaux
      //222222222222222222222222222222222222222222222222222222222222222222222222222222222222222
      cl_kernel ker2 = clCreateKernel(prg1, "matrix_cho2", NULL);
      
      // A priori, la matrice A est toujours enregistrée sur le GPU.
      clEnqueueWriteBuffer(cq, bufA, 0, 0, size, matA, 0, NULL, &ev_writeA);
      
      
      clSetKernelArg(ker2, 0, sizeof(bufA), &bufA);
      clSetKernelArg(ker2, 0, sizeof(int), &indice);
      clSetKernelArg(ker2, 0, sizeof(int),&taille_bloc);
      clSetKernelArg(ker2, 0, sizeof(int),&taille_mat);
      
      size_t globalDim[] = {TAILLE_MAT-i*TAILLE_BLOC,TAILLE_BLOC};//i correspond à un indice de boucle--> 1 à TAILLE_MAT/TAILLE_BLOC
      size_t localDim[] = {TAILLE_BLOC,TAILLE_BLOC};
      cl_event deps[] = {ev_writeA};
      cl_uint ret = clEnqueueNDRangeKernel(cq, ker2, 2, NULL, globalDim, localDim, 1, deps, &ev_ker2);
      printf("erreur : %d\n",ret);
      
      
      //Kernel : Mise à jour de tout le reste de la matrice
      //3333333333333333333333333333333333333333333333333333333333333333333333333333333333333333
      
      
      //TODO
      
      
      //------------------------------------------------------------------------
    }//fin de la boucle
  //clEnqueueReadBuffer(cq, bufA, 0, 0, size, matA, 1, &ev_ker2, &ev_readA);
  clFinish(cq);

  for (i=0; i<N; i++) {
    int j;
    for (j=0; j<N; j++) {
      printf("%f\t", matA[i*N + j]);
    }
    printf("\n");
  }

  /* for (i=0; i<15; i++) { */
  /*   printf("%f %f %f\t", matA[i],matB[i],matC[i]); */
  //}

  printf("\n");
  printf("Fini.\n");

  return 0;
}

